|
|
Главная
Статьи
Ссылки
Скачать
Скриншоты
Юмор
Почитать
Tools
Проекты
Обо мне
Гостевая
Форум
|
И OpenGL и CUDA работают с данными, находящимися в памяти GPU. Поэтому очень удобно иметь механизм, позволяющий им использовать эти данные совместно, без необходимости копирования в память CPU. CUDA предоставляет подобный механизм - CUDA OpenGL interoperability. Начиная с версии 3.0 в CUDA этот механизм взаимодействия с OpenGL был изменен и расширен.
Ключевым понятием в этом механизме является понятие графического ресурса, который может быть отображен в память CUDA, и обрабатываться средствами CUDA.
CUDA позволяет отображать в свое адресное пространство следующие объекты OpenGL - буфера, текстуры и рендербуферы. При этом буфера OpenGL отображаются в глобальную память, а текстуры и рендербуферы - в cudaArray.
Первым шагом взаимодействия CUDA и OpenGL является вызов функции cudaGLSetDevice (который заменяет вызов cudaSetDevice).
cudaError_t cudaGLSetGLDevice ( int device );
Прежде чем графический ресурс может быть использован в CUDA его необходимо зарегистрировать. Обратите внимание, что регистрация - это дорогостоящая операция, поэтому лучше всего для каждого ресурса выполнить ее всего один раз - в начале выполнения программы.
Для регистрации текстуры и рендербуфера используется функция cudaGLRegisterImage
cudaError_t cudaGraphicsGLRegisterImage ( struct cudaGraphicsResource ** resource, GLuint image, GLenum target,
unsigned int flags );
Зарегистрированную текстуру можно отобразить в cudaArray cudaGraphicsSubResourceGetMappedArray:
cudaError_t cudaGraphicsSubResourceGetMappedArray ( cudaArray ** array, struct cudaGraphicsResource * resource,
unsigned int arrayIndex, unsigned int mipLevel );
Обратите внимание, что пока CUDA поддерживает только текстуры с floating-point-компонентами (например, GL_RGBA_FLOAT32) и текстуры с ненормализованными целыми компонентами (например, GLRGBA8UI). Стандартные нормализованные целочисленные текстуры ( такие как GL_TGBA8) пока не поддерживаются.
Буфера OpenGL регистрируются при помощи функции cudaGraphicsGLRegisterBuffer:
cudaError_t cudaGraphicsGLRegisterBuffer ( struct cudaGraphicsResource ** resource, GLuint buffer, unsigned int flags );
Буфера отображаются в глобальную память CUDA.
При регистрации ресурса можно при помощи параметра flags сообщить CUDA о том, как именно будет использоваться данный ресурс, что может сделать работу с ним более эффективной. На данный момент поддерживаются следующие типы флагов:
После завершения всей работы с ресурсом, его следует "разрегистрировать" при помощи функции cudaGraphicsUnregisterResource:
cudaError_t cudaGraphicsUnregisterResource ( struct cudaGraphicsResource * resource );
Для непосредственной работы с уже зарегистрированным ресурсом необходимо сначала вызвать функцию cudaGraphicsMapResources, для отображения ресурса в память CUDA. После этого можно получить указатель для непосредственной работы с данными ресурса. Когда CUDA закончит работу с ресурсом, необходимо завершить отображение ресурса в память CUDA при помощи вызова функции cudaGraphicsUnmapResources.
Обратите внимание, что пока ресурс отображен в память CUDA доступ к нему из OpenGL невозможен (точнее дает непредсказуемый результат).
cudaError_t cudaGraphicsMapResources ( int count, struct cudaGraphicsResource ** resources, cudaStream_t stream ); cudaError_t cudaGraphicsUnmapResources ( int count, struct cudaGraphicsResource ** resources, cudaStream_t stream );
Отобразив ресурс в адресное пространство CUDA следует использовать функции cudaGraphicsResourceGetMappedPointer и cudaGraphicsSubResourceGetMappedArray для получения указателя, с которым может работать CUDA.
cudaError_t cudaGraphicsResourceGetMappedPointer ( void ** devPtr,
size_t * size, struct cudaGraphicsResource * resource );
cudaError_t cudaGraphicsSubResourceGetMappedArray ( cudaArray ** array, struct cudaGraphicsResource * resource,
unsigned int arrayIndex, unsigned int mipLevel );
Параметр arrayIndex служит для выбора грани для кубических текстур и текстурных массивов. Параметр mipLevel позволяет задать конкретный уровень в пирамиде mipmap-текстур.
Для упрощения дальнейшей работы с ресурсами OpenGL удобно сразу "завернуть" их в классы - CudaGlBuffer и CudaGlImage. Эти классы полностью инкапсулируют всю работу с ресурсами, включая регистрация/дерегистрацию (в конструкторе и деструкторе). Для непосредственного отображения ресурса в адресное пространство CUDA служат методы mapResource unmapResource.
После вызова mapResource для непосредственного доступа к данным ресурса их CUDA служат методы mappedPointer (для буферов OpenGL) и mappedArray (для текстур и рендербуферов OpenGL).
#ifndef __CUDA__GL_RESOURCE__
#define __CUDA__GL_RESOURCE__
#include <cuda.h>
#include <cuda_runtime_api.h>
#include <cuda_gl_interop.h>
#include <GL/gl.h>
#include "VertexBuffer.h"
class CudaGlBuffer // VBO
{
cudaGraphicsResource * resource;
VertexBuffer * buffer;
GLenum target;
public:
CudaGlBuffer ( VertexBuffer * buf, GLenum theTarget, unsigned int flags = cudaGraphicsMapFlagsWriteDiscard ) // cudaGraphicsMapFlagsReadOnly
{
buffer = buf;
target = theTarget;
buffer -> bind ( target );
cudaGraphicsGLRegisterBuffer ( &resource, buffer -> getId (), flags );
buffer -> unbind ();
}
~CudaGlBuffer ()
{
cudaGraphicsUnregisterResource ( resource );
}
bool mapResource ( cudaStream_t stream = 0 )
{
return cudaGraphicsMapResources ( 1, &resource, stream ) == cudaSuccess;
}
bool unmapResource ( cudaStream_t stream = 0 )
{
return cudaGraphicsUnmapResources ( 1, &resource, stream ) == cudaSuccess;
}
void * mappedPointer ( size_t& numBytes ) const
{
void * ptr;
if ( cudaGraphicsResourceGetMappedPointer ( &ptr, &numBytes, resource ) != cudaSuccess )
return NULL;
return ptr;
}
GLuint getId () const
{
return buffer -> getId ();
}
GLenum getTarget () const
{
return target;
}
cudaGraphicsResource * getResource () const
{
return resource;
}
};
class CudaGlImage // texture or renderbuffer
{
GLuint image;
GLenum target;
cudaGraphicsResource * resource;
public:
CudaGlImage ( GLuint theImage, GLenum theTarget, unsigned int flags = cudaGraphicsMapFlagsWriteDiscard ) // cudaGraphicsMapFlagsReadOnly, cudaGraphicsMapFlagsNone
{
image = theImage;
target = theTarget;
cudaGraphicsGLRegisterImage ( &resource, image, target, flags );
}
~CudaGlImage ()
{
cudaGraphicsUnregisterResource ( resource );
}
bool mapResource ( cudaStream_t stream = 0 )
{
return cudaGraphicsMapResources ( 1, &resource, stream ) == cudaSuccess;
}
bool unmapResource ( cudaStream_t stream = 0 )
{
return cudaGraphicsUnmapResources ( 1, &resource, stream ) == cudaSuccess;
}
cudaArray * mappedArray ( unsigned int index = 0, unsigned int mipLevel = 0 ) const
{
cudaArray * array;
if ( cudaGraphicsSubResourceGetMappedArray ( &array, resource, index, mipLevel ) != cudaSuccess )
return NULL;
return array;
}
GLuint getImage () const
{
return image;
}
GLenum getTarget () const
{
return target;
}
};
#endif
Далее мы рассмотрим несколько примеров взаимодействия CUDA и OpenGL. Простейшим из этих примеров будет создание и рендеринг анимированной водной поверхности. Координаты вершин и значения нормали в них для каждого кадра будем рассчитывать при помощи ядра CUDA, а непосредственно рендеринг осуществлять средствами OpenGL.
Нам понадобятся два вершинных буфера (VBO) - один для хранения вершин, второй - для хранения нормалей. На следующем листинге приводится само ядро и вызывающая его функция:
#define TEX_SIZE 512
#define k1 2.345431f
#define k2 1.12312312f
#define k3 3.784728394f
#define scale 0.03f
__global__ void animateKernel ( float4 * pos, float4 * norm, int w, int wh, int hh, float time )
{
int ix = blockIdx.x * blockDim.x + threadIdx.x;
int iy = blockIdx.y * blockDim.y + threadIdx.y;
int i = iy * w + ix;
float x = (ix - wh) * 0.03f;
float y = (iy - hh) * 0.03f;
float v1 = x * 0.47f + y * 1.3f + k1 * time;
float v2 = x * 2.9f + y * 3.1f + k2 * time;
float v3 = x * 5.577f - y * 7.57f + k3 * time;
float nx = scale * ( 2.77f * 0.47f * __cosf ( v1 ) + 2.9f * __cosf ( v2 ) + 0.357f * 5.577f * __cosf ( v3 ) );
float ny = scale * ( 2.77f * 1.3f * __cosf ( v1 ) + 3.1f * __cosf ( v2 ) - 0.357f * 7.57f * __cosf ( v3 ) );
float nz = 1.0f / sqrtf ( nx * nx + ny * ny );
pos [i] = make_float4 ( x, y, scale * ( 2.77f * __sinf ( v1 ) + __sinf ( v2 ) + 0.357f * __sinf ( v3 ) ), 1.0f );
norm [i] = make_float4 ( nx, ny, nz, 0.0f );
}
extern "C" void buildVertices ( float4 * v, float4 * n, int w, int h, float time )
{
dim3 threads ( 16, 16 );
dim3 blocks ( w / 16, h / 16 );
animateKernel <<<blocks, threads>>> ( v, n, w, w / 2, h / 2, time );
cudaThreadSynchronize ();
}
Ядро использует аналитическое задание поверхности воды при помощи суммы нескольких гармоник. Вектор нормали также вычисляется аналитически.
Ниже приводится основной код на С++, осуществляющий рендеринг анимированной поверхности воды. Обратите внимание, что ряд постоянно используемых во многих примерах функций удалены из листинга (полный текст всех программ доступен по ссылке в конце статьи).
#include "libExt.h"
#ifdef MACOSX
#include <GLUT/glut.h>
#else
#include <glut.h>
#endif
#include <stdio.h>
#include <stdlib.h>
#include "libTexture.h"
#include "TypeDefs.h"
#include "Vector3D.h"
#include "Vector2D.h"
#include "boxes.h"
#include "GlslProgram.h"
#include "CudaGlResource.h"
#define TEX_SIZE 512 // width & height of texture
Vector3D eye ( -10, -10, 5 ); // camera position
Vector3D light ( 5, 0, 4 ); // light position
float angle = 0;
Vector3D rot ( 0, 0, 0 );
int mouseOldX = 0;
int mouseOldY = 0;
float waveScale = 0.2;
float time = 0.0f;
unsigned reflectionMap; // cubic map with reflection
CudaGlBuffer * pos = NULL;
CudaGlBuffer * norm = NULL;
VertexBuffer * vertexBuffer; // vertex coordinates
VertexBuffer * normalBuffer; // normal buffer
VertexBuffer * indexBuffer; // vertex index buffer
GlslProgram program;
extern "C" void buildVertices ( float4 * v, float4 * n, int w, int h, float time );
void createIndexBuffer ()
{
int * data = (int *) malloc ( (TEX_SIZE-1)*(TEX_SIZE-1)*6*sizeof (int) );
int k = 0;
for ( int i = 0; i < TEX_SIZE - 1; i++ )
for ( int j = 0; j < TEX_SIZE - 1; j++ )
{
data [k] = i + TEX_SIZE*j; // first triangle (i,j)-(i+1,j) - (i,j+1)
data [k+1] = i + TEX_SIZE*j + 1;
data [k+2] = i + TEX_SIZE*j + TEX_SIZE;
data [k+3] = i + TEX_SIZE*j + 1; // second triangle (i+1,j)-(i+1,j+1) - (i,j+1)
data [k+4] = i + TEX_SIZE*j + TEX_SIZE + 1;
data [k+5] = i + TEX_SIZE*j + TEX_SIZE;
k += 6;
}
indexBuffer = new VertexBuffer ();
indexBuffer -> bind ( GL_ELEMENT_ARRAY_BUFFER_ARB );
indexBuffer -> setData ( (TEX_SIZE - 1) * (TEX_SIZE - 1) * 6 * sizeof ( int ), data, GL_STATIC_DRAW );
indexBuffer -> unbind ();
free ( data );
}
void createVertexBuffer ()
{
vertexBuffer = new VertexBuffer ();
vertexBuffer -> bind ( GL_PIXEL_PACK_BUFFER_ARB );
vertexBuffer -> setData ( TEX_SIZE * TEX_SIZE * 4 * sizeof ( float ), NULL, GL_DYNAMIC_DRAW );
vertexBuffer -> unbind ();
}
void createNormalBuffer ()
{
normalBuffer = new VertexBuffer ();
normalBuffer -> bind ( GL_PIXEL_PACK_BUFFER_ARB );
normalBuffer -> setData ( TEX_SIZE * TEX_SIZE * 4 * sizeof ( float ), NULL, GL_DYNAMIC_DRAW );
normalBuffer -> unbind ();
}
void display ()
{
float scale = 1.0f;
size_t vSize, nSize;
pos -> mapResource ();
norm -> mapResource ();
buildVertices ( (float4 *) pos -> mappedPointer ( vSize ),
(float4 *) norm -> mappedPointer ( nSize ), TEX_SIZE, TEX_SIZE, time );
pos -> unmapResource ();
norm -> unmapResource ();
glClear ( GL_COLOR_BUFFER_BIT | GL_DEPTH_BUFFER_BIT );
glMatrixMode ( GL_MODELVIEW );
glPushMatrix ();
glRotatef ( rot.x, 1, 0, 0 );
glRotatef ( rot.y, 0, 1, 0 );
glRotatef ( rot.z, 0, 0, 1 );
glScalef ( scale, scale, 2 );
glActiveTextureARB ( GL_TEXTURE0_ARB );
glBindTexture ( GL_TEXTURE_CUBE_MAP, reflectionMap );
glPushClientAttrib ( GL_CLIENT_VERTEX_ARRAY_BIT );
glEnableClientState ( GL_VERTEX_ARRAY );
vertexBuffer -> bind ( GL_ARRAY_BUFFER_ARB );
glVertexPointer ( 4, GL_FLOAT, 0, NULL );
glEnableClientState ( GL_TEXTURE_COORD_ARRAY );
normalBuffer -> bind ( GL_ARRAY_BUFFER_ARB );
glTexCoordPointer ( 4, GL_FLOAT, 0, NULL );
indexBuffer -> bind ( GL_ELEMENT_ARRAY_BUFFER_ARB );
glIndexPointer ( GL_UNSIGNED_INT, 0, 0 );
program.bind ();
glDrawElements ( GL_TRIANGLES, (TEX_SIZE-1)*(TEX_SIZE-1)*6, GL_UNSIGNED_INT, 0 );
program.unbind ();
glPopClientAttrib ();
glPopMatrix ();
glutSwapBuffers ();
}
int main ( int argc, char * argv [] )
{
// initialize glut
glutInit ( &argc, argv );
glutInitDisplayMode ( GLUT_DOUBLE | GLUT_RGBA | GLUT_DEPTH );
glutInitWindowSize ( 640, 480 );
// create window
glutCreateWindow ( "OpenGL/CUDA interoperability: vertex buffers" );
// register handlers
glutDisplayFunc ( display );
glutReshapeFunc ( reshape );
glutKeyboardFunc ( key );
glutMouseFunc ( mouse );
glutMotionFunc ( motion );
glutIdleFunc ( animate );
init ();
initExtensions ();
if ( !GlslProgram :: isSupported () )
{
printf ( "GLSL not supported\n" );
return 1;
}
const char * faces [6] =
{
"../../Textures/Cubemaps/skyrt.bmp",
"../../Textures/Cubemaps/skylf.bmp",
"../../Textures/Cubemaps/skydn.bmp",
"../../Textures/Cubemaps/skyup.bmp",
"../../Textures/Cubemaps/skybk.bmp",
"../../Textures/Cubemaps/skyfr.bmp",
};
if ( !program.loadShaders ( "water-draw.vsh", "water-draw.fsh" ) )
{
printf ( "Error loading water-draw shaders:\n%s\n", program.getLog ().c_str () );
return 3;
}
program.bind ();
program.setTexture ( "reflectionMap", 0 );
program.setUniformVector ( "eye", eye );
program.unbind ();
reflectionMap = createCubeMap ( true, faces );
// initialize CUDA device for OpenGL
cudaGLSetGLDevice ( 0 );
createVertexBuffer ();
createNormalBuffer ();
createIndexBuffer ();
pos = new CudaGlBuffer ( vertexBuffer, GL_ARRAY_BUFFER );
norm = new CudaGlBuffer ( normalBuffer, GL_ARRAY_BUFFER );
glutMainLoop ();
return 0;
}
Далее мы рассмотрим image processing средствами CUDA. При этом для записи результатов в текстуру возможны несколько путей. Можно осуществить запись результатов ядром в обычный массив в глобальной памяти, а затем уже скопировать данные в cudaArray, соответствующий выходной текстуре.
Второй вариант заключается в использовании расширения ARB_pixel_buffer_object, позволяющего копировать данные из вершинного буфера прямо в текстуру. В этом случае создается специальный вершинный буфер, который отображается в глобальную память CUDA. Ядро записывает в него результат работы. Далее через расширение ARB_pixel_buffer_object производится запись данных из этого вершинного буфера в текстуру.
Ниже приводится соответствующий код на CUDA для первого случая - вся работа идет через cudaArray. Поскольку ядро не может напрямую писать в cudaArray, полученный отображением текстуры, то ядро выводит результат в заранее выделенный блок глобальной памяти, а затем при помощи функции cudaMemcpyToArray, копирует результат в cudaArray, соответствующий выходной текстуре.
//
//
// CUDA kernal to convert image to grey-scale , keep intact one specific color hue
//
#define EPS 0.006f
typedef unsigned char byte;
texture<uchar4, 2, cudaReadModeElementType> inTex;
__device__ float3 rgbToHsv ( const float3 c )
{
float mn = min ( min ( c.x, c.y ), c.z );
float mx = max ( max ( c.x, c.y ), c.z );
float delta = mx - mn;
float h, s;
if ( mx > 0.001f )
{
s = delta / mx;
if ( c.x == mx )
h = ( c.y - c.z ) / delta;
else
if ( c.y == mx )
h = 2.0f + ( c.z - c.x ) / delta;
else
h = 4.0f + ( c.x - c.y ) / delta;
}
else
{
s = 0.0f;
h = 0.0f;
}
return make_float3 ( h / 6.0f, s, mx );
}
__global__ void ppKernel ( uchar4 * out, int w, int h )
{
int x = blockIdx.x * blockDim.x + threadIdx.x;
int y = blockIdx.y * blockDim.y + threadIdx.y;
int i = y * w + x;
uchar4 res = tex2D ( inTex, x, y );
float scale = 1.0f / 255.0f;
float3 clr = make_float3 ( res.x * scale, res.y * scale, res.z * scale );
float3 hsv = rgbToHsv ( clr ); // convert to HSV
float l = 0.3f * clr.x + 0.59f * clr.y + 0.11f * clr.z; // intensity of a color
if ( abs ( hsv.x - 0.077f ) < EPS ) // hit hue with enough precision => keep color
out [i] = res;
else // use greyscale image
{
byte r = (byte)( 255.0f * l);
out [i] = make_uchar4 ( r, r, r, 255 );
}
}
extern "C" void doPostprocess ( cudaArray * inMap, cudaArray * outMap, uchar4 * buf, int pitch, int w, int h )
{
cudaBindTextureToArray ( inTex, inMap );
dim3 threads ( 16, 16 );
dim3 blocks ( w / threads.x, h / threads.y);
ppKernel <<<blocks, threads>>> ( buf, w, h );
cudaMemcpyToArray ( outMap, 0, 0, buf, w * h * sizeof(uchar4), cudaMemcpyDeviceToDevice );
cudaThreadSynchronize ();
cudaUnbindTexture ( inTex );
}
На следующем листинге приводится соответствующий код на С++. Обратите внимание, что поскольку мы использование специальных шейдеров, служащий для записи и чтения в текстуры с ненормализованным целочисленным форматом компонент.
#include "libExt.h"
#ifdef MACOSX
#include <GLUT/glut.h>
#else
#include <glut.h>
#endif
#include <stdio.h>
#include <stdlib.h>
#include "libTexture.h"
#include "TypeDefs.h"
#include "Vector3D.h"
#include "Vector2D.h"
#include "boxes.h"
#include "FrameBuffer.h"
#include "GlslProgram.h"
#include "CudaGlResource.h"
Vector3D eye ( -0.5, -0.5, 1.5 ); // camera position
unsigned decalMap; // decal (diffuse) texture
unsigned stoneMap;
unsigned teapotMap;
unsigned screenMap;
unsigned outMap;
float angle = 0;
float rot = 0;
bool useFilter = true;
CudaGlImage * src = NULL;
CudaGlImage * dst = NULL;
uchar4 * buf = NULL;
FrameBuffer buffer ( 640, 480, FrameBuffer :: depth32 );
GlslProgram program1; // render to uchar4 texture
GlslProgram program2; // render to normalized color buffer
void renderToBuffer ();
void postProcess ();
extern "C" void doPostprocess ( cudaArray * inArray, cudaArray * outArray, uchar4 * buf, int pitch, int w, int h );
void displayBoxes ()
{
glMatrixMode ( GL_MODELVIEW );
glPushMatrix ();
glRotatef ( rot, 0, 0, 1 );
drawBox ( Vector3D ( -5, -5, 0 ), Vector3D ( 10, 10, 3 ), stoneMap, false );
drawBox ( Vector3D ( 3, 2, 0.5 ), Vector3D ( 1, 2, 2 ), decalMap );
glBindTexture ( GL_TEXTURE_2D, teapotMap );
glTranslatef ( 0.2, 1, 1.5 );
glRotatef ( angle * 45.3, 1, 0, 0 );
glRotatef ( angle * 57.2, 0, 1, 0 );
glutSolidTeapot ( 0.3 );
glPopMatrix ();
}
void display ()
{
renderToBuffer ();
postProcess ();
glClear ( GL_COLOR_BUFFER_BIT | GL_DEPTH_BUFFER_BIT );
startOrtho ();
program2.bind ();
glEnable ( GL_TEXTURE_2D );
glBindTexture ( GL_TEXTURE_2D, outMap );
glBegin ( GL_QUADS );
glTexCoord2f ( 0, 0 );
glVertex2f ( 0, 0 );
glTexCoord2f ( 1, 0 );
glVertex2f ( buffer.getWidth (), 0 );
glTexCoord2f ( 1, 1 );
glVertex2f ( buffer.getWidth (), buffer.getHeight () );
glTexCoord2f ( 0, 1 );
glVertex2f ( 0, buffer.getHeight () );
glEnd ();
glDisable ( GL_TEXTURE_2D );
glBindTexture ( GL_TEXTURE_2D, 0 );
program2.unbind ();
endOrtho ();
glutSwapBuffers ();
}
void renderToBuffer ()
{
glBindTexture ( GL_TEXTURE_2D, 0 );
program1.bind ();
buffer.bind ();
glClearColor ( 0, 0, 0, 1 );
glClear ( GL_COLOR_BUFFER_BIT | GL_DEPTH_BUFFER_BIT );
reshape ( buffer.getWidth (), buffer.getHeight () );
displayBoxes ();
buffer.unbind ();
program1.unbind ();
}
void postProcess ()
{
src -> mapResource ();
dst -> mapResource ();
cudaArray * inArray = src -> mappedArray ();
cudaArray * outArray = dst -> mappedArray ();
doPostprocess ( inArray, outArray, buf, 0, buffer.getWidth (), buffer.getHeight () );
src -> unmapResource ();
dst -> unmapResource ();
}
GLuint createTextureDst ( int w, int h )
{
GLuint tex;
glGenTextures ( 1, &tex );
glBindTexture ( GL_TEXTURE_2D, tex );
glTexParameteri ( GL_TEXTURE_2D, GL_TEXTURE_WRAP_S, GL_CLAMP_TO_EDGE );
glTexParameteri ( GL_TEXTURE_2D, GL_TEXTURE_WRAP_T, GL_CLAMP_TO_EDGE );
glTexParameteri ( GL_TEXTURE_2D, GL_TEXTURE_MIN_FILTER, GL_NEAREST );
glTexParameteri ( GL_TEXTURE_2D, GL_TEXTURE_MAG_FILTER, GL_NEAREST );
glTexImage2D ( GL_TEXTURE_2D, 0, GL_RGBA8UI_EXT, w, h, 0, GL_RGBA_INTEGER_EXT, GL_UNSIGNED_BYTE, NULL );
return tex;
}
int main ( int argc, char * argv [] )
{
// initialize glut
glutInit ( &argc, argv );
glutInitDisplayMode ( GLUT_DOUBLE | GLUT_RGBA | GLUT_DEPTH );
glutInitWindowSize ( buffer.getWidth (), buffer.getHeight () );
// create window
glutCreateWindow ( "OpenGL/CUDA interoperability: postprocessing" );
// register handlers
glutDisplayFunc ( display );
glutReshapeFunc ( reshape );
glutKeyboardFunc ( key );
glutSpecialFunc ( specialKey );
glutIdleFunc ( animate );
init ();
initExtensions ();
assertExtensionsSupported ( "GL_ARB_shading_language_100 GL_ARB_shader_objects EXT_framebuffer_object ARB_texture_rectangle" );
outMap = createTextureDst ( buffer.getWidth (), buffer.getHeight () );
decalMap = createTexture2D ( true, "../../Textures/oak.bmp" );
stoneMap = createTexture2D ( true, "../../Textures/block.bmp" );
teapotMap = createTexture2D ( true, "../../Textures/Oxidated.jpg" );
screenMap = buffer.createColorTexture ( GL_RGBA_INTEGER_EXT, GL_RGBA8UI_EXT );
buffer.create ();
buffer.bind ();
buffer.attachColorTexture ( GL_TEXTURE_2D, screenMap );
if ( !buffer.isOk () )
printf ( "Error with framebuffer\n" );
buffer.unbind ();
if ( !program1.loadShaders ( "draw.vsh", "draw-uint.fsh" ) )
{
printf ( "Error loading draw-uint shaders:\n%s\n", program1.getLog ().c_str () );
return 3;
}
program1.bind ();
program1.setTexture ( "texImage", 0 );
program1.unbind ();
if ( !program2.loadShaders ( "draw.vsh", "draw-color.fsh" ) )
{
printf ( "Error loading draw-uint shaders:\n%s\n", program2.getLog ().c_str () );
return 3;
}
program2.bind ();
program2.setTexture ( "texImage", 0 );
program2.unbind ();
// initialize CUDA device for OpenGL
cudaGLSetGLDevice ( 0 );
cudaMalloc ( (void **) &buf, buffer.getWidth () * buffer.getHeight () * 4 );
src = new CudaGlImage ( buffer.getColorBuffer (), GL_TEXTURE_2D, cudaGraphicsMapFlagsNone );
dst = new CudaGlImage ( outMap, GL_TEXTURE_2D, cudaGraphicsMapFlagsNone );
glutMainLoop ();
return 0;
}
В случае использования расширения ARB_pixel_buffer_object код на CUDA получается даже проще - ядро сразу записывает результат в глобальную память, только это уже не заранее выделенный блок глобальной памяти, а отображенный вершинный буфер.
//
// CUDA kernal to convert image to grey-scale , keep intact one specific color hue
//
#define EPS 0.006f
typedef unsigned char byte;
texture<uchar4, 2, cudaReadModeElementType> inTex;
__device__ float3 rgbToHsv ( const float3 c )
{
float mn = min ( min ( c.x, c.y ), c.z );
float mx = max ( max ( c.x, c.y ), c.z );
float delta = mx - mn;
float h, s;
if ( mx > 0.001f )
{
s = delta / mx;
if ( c.x == mx )
h = ( c.y - c.z ) / delta;
else
if ( c.y == mx )
h = 2.0f + ( c.z - c.x ) / delta;
else
h = 4.0f + ( c.x - c.y ) / delta;
}
else
{
s = 0.0f;
h = 0.0f;
}
return make_float3 ( h / 6.0f, s, mx );
}
__global__ void ppKernel ( int w, int h, uchar4 * outPtr )
{
int x = blockIdx.x * blockDim.x + threadIdx.x;
int y = blockIdx.y * blockDim.y + threadIdx.y;
int i = y * w + x;
uchar4 res = tex2D ( inTex, x, y );
float scale = 1.0f / 255.0f;
float3 clr = make_float3 ( res.x * scale, res.y * scale, res.z * scale );
float3 hsv = rgbToHsv ( clr ); // convert to HSV
float l = 0.3f * clr.x + 0.59f * clr.y + 0.11f * clr.z; // intensity of a color
if ( abs ( hsv.x - 0.077f ) < EPS ) // hit hue with enough precision => keep color
outPtr [i] = res;
else // use greyscale image
{
byte r = (byte)(255.0f * l);
outPtr [i] = make_uchar4 ( r, r, r, 255 ); //make_uchar4 ( x, y, 100, 100 );
}
}
extern "C" void doPostprocess ( cudaArray * inMap, uchar4 * outPtr, int w, int h )
{
cudaBindTextureToArray ( inTex, inMap );
dim3 threads ( 16, 16 );
dim3 blocks ( w / threads.x, h / threads.y);
ppKernel <<<blocks, threads>>> ( w, h, outPtr );
cudaThreadSynchronize ();
cudaUnbindTexture ( inTex );
}
Ниже приводится соответствующий кода на С++.
#include "libExt.h"
#ifdef MACOSX
#include <GLUT/glut.h>
#else
#include <glut.h>
#endif
#include <stdio.h>
#include <stdlib.h>
#include "libTexture.h"
#include "TypeDefs.h"
#include "Vector3D.h"
#include "Vector2D.h"
#include "boxes.h"
#include "FrameBuffer.h"
#include "GlslProgram.h"
#include "CudaGlResource.h"
Vector3D eye ( -0.5, -0.5, 1.5 ); // camera position
unsigned decalMap; // decal (diffuse) texture
unsigned stoneMap;
unsigned teapotMap;
unsigned screenMap;
unsigned outMap;
float angle = 0;
float rot = 0;
bool useFilter = true;
CudaGlImage * src = NULL;
CudaGlBuffer * out = NULL;
VertexBuffer * imageBuffer; // vertex coordinates
FrameBuffer buffer ( 640, 480, FrameBuffer :: depth32 );
GlslProgram program1;
void renderToBuffer ();
void postProcess ();
extern "C" void doPostprocess ( cudaArray * inArray, uchar4 * outPtr, int w, int h );
void createVertexBuffer ()
{
imageBuffer = new VertexBuffer ();
imageBuffer -> bind ( GL_PIXEL_PACK_BUFFER_ARB );
imageBuffer -> setData ( buffer.getWidth () * buffer.getHeight () * 4, NULL, GL_DYNAMIC_DRAW );
imageBuffer -> unbind ();
}
void displayBoxes ()
{
glMatrixMode ( GL_MODELVIEW );
glPushMatrix ();
glRotatef ( rot, 0, 0, 1 );
drawBox ( Vector3D ( -5, -5, 0 ), Vector3D ( 10, 10, 3 ), stoneMap, false );
drawBox ( Vector3D ( 3, 2, 0.5 ), Vector3D ( 1, 2, 2 ), decalMap );
glBindTexture ( GL_TEXTURE_2D, teapotMap );
glTranslatef ( 0.2, 1, 1.5 );
glRotatef ( angle * 45.3, 1, 0, 0 );
glRotatef ( angle * 57.2, 0, 1, 0 );
glutSolidTeapot ( 0.3 );
glPopMatrix ();
}
void display ()
{
renderToBuffer ();
postProcess ();
glClear ( GL_COLOR_BUFFER_BIT | GL_DEPTH_BUFFER_BIT );
startOrtho ();
glEnable ( GL_TEXTURE_RECTANGLE_ARB );
glBindTexture ( GL_TEXTURE_RECTANGLE_ARB, outMap );
glBegin ( GL_QUADS );
glTexCoord2f ( 0, 0 );
glVertex2f ( 0, 0 );
glTexCoord2f ( buffer.getWidth (), 0 );
glVertex2f ( buffer.getWidth (), 0 );
glTexCoord2f ( buffer.getWidth (), buffer.getHeight () );
glVertex2f ( buffer.getWidth (), buffer.getHeight () );
glTexCoord2f ( 0, buffer.getHeight () );
glVertex2f ( 0, buffer.getHeight () );
glEnd ();
glDisable ( GL_TEXTURE_RECTANGLE_ARB );
glBindTexture ( GL_TEXTURE_RECTANGLE_ARB, 0 );
endOrtho ();
glutSwapBuffers ();
}
void renderToBuffer ()
{
glBindTexture ( GL_TEXTURE_2D, 0 );
program1.bind ();
buffer.bind ();
glClearColor ( 0, 0, 0, 1 );
glClear ( GL_COLOR_BUFFER_BIT | GL_DEPTH_BUFFER_BIT );
reshape ( buffer.getWidth (), buffer.getHeight () );
displayBoxes ();
buffer.unbind ();
program1.unbind ();
}
void postProcess ()
{
src -> mapResource ();
out -> mapResource ();
size_t size;
cudaArray * inArray = src -> mappedArray ();
uchar4 * res = (uchar4 *)out -> mappedPointer ( size );
doPostprocess ( inArray, res, buffer.getWidth (), buffer.getHeight () );
src -> unmapResource ();
out -> unmapResource ();
imageBuffer -> bind ( GL_PIXEL_UNPACK_BUFFER_ARB );
glBindTexture ( GL_TEXTURE_RECTANGLE_ARB, outMap );
glTexSubImage2D ( GL_TEXTURE_RECTANGLE_ARB, 0, 0, 0, buffer.getWidth (), buffer.getHeight (), GL_RGBA, GL_UNSIGNED_BYTE, NULL );
imageBuffer -> unbind ();
glBindBufferARB ( GL_PIXEL_UNPACK_BUFFER_ARB, 0 );
glBindBufferARB ( GL_PIXEL_PACK_BUFFER_ARB, 0 );
}
GLuint createTextureDst ( int w, int h )
{
GLuint tex;
glGenTextures ( 1, &tex );
glBindTexture ( GL_TEXTURE_RECTANGLE_ARB, tex );
glTexParameteri ( GL_TEXTURE_RECTANGLE_ARB, GL_TEXTURE_WRAP_S, GL_CLAMP_TO_EDGE );
glTexParameteri ( GL_TEXTURE_RECTANGLE_ARB, GL_TEXTURE_WRAP_T, GL_CLAMP_TO_EDGE );
glTexParameteri ( GL_TEXTURE_RECTANGLE_ARB, GL_TEXTURE_MIN_FILTER, GL_NEAREST );
glTexParameteri ( GL_TEXTURE_RECTANGLE_ARB, GL_TEXTURE_MAG_FILTER, GL_NEAREST );
glTexImage2D ( GL_TEXTURE_RECTANGLE_ARB, 0, GL_RGBA8, w, h, 0, GL_RGBA, GL_UNSIGNED_BYTE, NULL );
glBindTexture ( GL_TEXTURE_RECTANGLE_ARB, 0 );
return tex;
}
int main ( int argc, char * argv [] )
{
// initialize glut
glutInit ( &argc, argv );
glutInitDisplayMode ( GLUT_DOUBLE | GLUT_RGBA | GLUT_DEPTH );
glutInitWindowSize ( buffer.getWidth (), buffer.getHeight () );
// create window
glutCreateWindow ( "OpenGL/CUDA interoperability: postprocessing" );
// register handlers
glutDisplayFunc ( display );
glutReshapeFunc ( reshape );
glutKeyboardFunc ( key );
glutSpecialFunc ( specialKey );
glutIdleFunc ( animate );
init ();
initExtensions ();
decalMap = createTexture2D ( true, "../../Textures/oak.bmp" );
stoneMap = createTexture2D ( true, "../../Textures/block.bmp" );
teapotMap = createTexture2D ( true, "../../Textures/Oxidated.jpg" );
outMap = createTextureDst ( buffer.getWidth (), buffer.getHeight () );
screenMap = buffer.createColorTexture ( GL_RGBA );
buffer.create ();
buffer.bind ();
buffer.attachColorTexture ( GL_TEXTURE_2D, screenMap );
if ( !buffer.isOk () )
printf ( "Error with framebuffer\n" );
buffer.unbind ();
if ( !program1.loadShaders ( "draw.vsh", "draw.fsh" ) )
{
printf ( "Error loading draw shaders:\n%s\n", program1.getLog ().c_str () );
return 3;
}
program1.bind ();
program1.setTexture ( "texImage", 0 );
program1.unbind ();
// initialize CUDA device for OpenGL
cudaGLSetGLDevice ( 0 );
src = new CudaGlImage ( buffer.getColorBuffer (), GL_TEXTURE_2D, cudaGraphicsMapFlagsNone );
createVertexBuffer ();
out = new CudaGlBuffer ( imageBuffer, GL_ARRAY_BUFFER );
glutMainLoop ();
return 0;
}
По этой ссылке можно скачать весь исходный код к этой статье. Также доступны для скачивания откомпилированные версии для M$ Windows.
